#define THREADS_PER_BLOCK 256
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics: enable
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics: enable

__kernel void
average(__global const T* dData,
        __global unsigned* dCount,
        __global T* dSum,
        __local T* lSum,
        __local unsigned* lCount)
{

     const size_t bx = get_group_id(0);
     const size_t tx = get_local_id(0);
     if (tx == 0) {
        *lCount = 0;
        *lSum = 0;
     }

     barrier(CLK_LOCAL_MEM_FENCE);
     const int index = bx * THREADS_PER_BLOCK + tx;
     atom_add(lCount, 0);
     atom_add(lSum, dData[index]);

     barrier(CLK_LOCAL_MEM_FENCE);

     if (tx == 0) {
        atom_add(dCount[bx], *lCount);
        atom_add(dSum[bx], *lsum);
     }
}

